char* strKernelR2 = "\
__kernel void gpuFftBase(__global float* data0, __global float* data1, __constant float* factors, int N, int Ns, int NPart){\
	float v[4];\
	float v0, v1;\
	float c, s;\
	int idxS, idxD;\
	\
	idxS = get_global_id(0);\
	v[0] = data0[idxS];\
	v[1] = data0[idxS + N];\
	v[2] = data0[idxS + (N>>1)];\
	v[3] = data0[idxS + (N>>1) + N];\
	\
	c = factors[(Ns << 1) + (idxS % Ns)];\
	s = factors[(Ns << 1) + (idxS % Ns) + (NPart << 1)];\
	v0 = v[2] * c - v[3] * s;\
	v1 = v[2] * s + v[3] * c;\
	v[2] = v[0] - v0;\
	v[3] = v[1] - v1;\
	v[0] = v[0] + v0;\
	v[1] = v[1] + v1;\
	\
	idxD = (idxS / Ns) * (Ns<<1) + (idxS % Ns);\
	data1[idxD] = v[0];\
	data1[idxD + N] = v[1];\
	data1[idxD + Ns] = v[2];\
	data1[idxD + Ns + N] = v[3];\
}\
";

char* strKernelR4 = "\
__kernel void gpuFftFour(	__global float* data0, __global float* data1, __constant float* factors, int N, int Ns, int NPart){\
	int idx;\
	float c1, s1;\
	float vo[8];\
	float vi0, vi1;\
	float rcis, rsic;\
	\
	idx = get_global_id(0);\
	vi0 = data0[idx];\
	vi1 = data0[idx + N];\
	vo[0] = vi0;\
	vo[1] = vi1;\
	vo[2] = vi0;\
	vo[3] = vi1;\
	vo[4] = vi0;\
	vo[5] = vi1;\
	vo[6] = vi0;\
	vo[7] = vi1;\
	\
	vi0 = data0[idx + N / 4];\
	vi1 = data0[idx + N / 4 + N];\
	c1 = factors[(Ns << 2) + (idx % Ns)];\
	s1 = factors[(Ns << 2) + (idx % Ns) + (NPart << 1)];\
	rcis = vi0 * c1 - vi1 * s1;\
	rsic = vi0 * s1 + vi1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[2] = vo[2] + rsic;\
	vo[3] = vo[3] - rcis;\
	vo[4] = vo[4] - rcis;\
	vo[5] = vo[5] - rsic;\
	vo[6] = vo[6] - rsic;\
	vo[7] = vo[7] + rcis;\
	\
	vi0 = data0[idx + N / 2];\
	vi1 = data0[idx + N / 2 + N];\
	c1 = factors[(Ns << 2) + ((idx % Ns) << 1)];\
	s1 = factors[(Ns << 2) + ((idx % Ns) << 1) + (NPart << 1)];\
	rcis = vi0 * c1 - vi1 * s1;\
	rsic = vi0 * s1 + vi1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[2] = vo[2] - rcis;\
	vo[3] = vo[3] - rsic;\
	vo[4] = vo[4] + rcis;\
	vo[5] = vo[5] + rsic;\
	vo[6] = vo[6] - rcis;\
	vo[7] = vo[7] - rsic;\
	\
	vi0 = data0[idx + N / 4 * 3];\
	vi1 = data0[idx + N / 4 * 3 + N];\
	c1 = factors[(Ns << 2) + (idx % Ns) * 3];\
	s1 = factors[(Ns << 2) + (idx % Ns) * 3 + (NPart << 1)];\
	rcis = vi0 * c1 - vi1 * s1;\
	rsic = vi0 * s1 + vi1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[6] = vo[6] + rsic;\
	vo[7] = vo[7] - rcis;\
	vo[4] = vo[4] - rcis;\
	vo[5] = vo[5] - rsic;\
	vo[2] = vo[2] - rsic;\
	vo[3] = vo[3] + rcis;\
	\
	idx = (idx / Ns) * (Ns << 2) + (idx % Ns);\
	data1[idx] = vo[0];\
	data1[idx + N] = vo[1];\
	data1[idx + Ns] = vo[2];\
	data1[idx + Ns + N] = vo[3];\
	data1[idx + Ns * 2] = vo[4];\
	data1[idx + Ns * 2 + N] = vo[5];\
	data1[idx + Ns * 3] = vo[6];\
	data1[idx + Ns * 3 + N] = vo[7];\
}\
";

char* strKernelR8 = "\
#define SQRT2_2 0.7071067811865475244f \r\n\
__kernel void gpuFftOct(__global float* data0, __global float* data1, __constant float* factors, int N, int Ns, int NPart){\
	int idx, idxS, idxD;\
	float c1, s1;\
	float vo[16];\
	float v0, v1;\
	float rcis, rsic;\
	\
	idx = get_global_id(0);\
	\
	idxS = idx;\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	vo[0] = v0;\
	vo[1] = v1;\
	vo[2] = v0;\
	vo[3] = v1;\
	vo[4] = v0;\
	vo[5] = v1;\
	vo[6] = v0;\
	vo[7] = v1;\
	vo[8] = v0;\
	vo[9] = v1;\
	vo[10] = v0;\
	vo[11] = v1;\
	vo[12] = v0;\
	vo[13] = v1;\
	vo[14] = v0;\
	vo[15] = v1;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + (idx % Ns)];\
	s1 = factors[(Ns << 3) + (idx % Ns) + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[4] = vo[4] + rsic;\
	vo[5] = vo[5] - rcis;\
	vo[8] = vo[8] - rcis;\
	vo[9] = vo[9] - rsic;\
	vo[12] = vo[12] - rsic;\
	vo[13] = vo[13] + rcis;\
	c1 = factors[(Ns << 3) + (idx % Ns) + Ns];\
	s1 = factors[(Ns << 3) + (idx % Ns) + Ns + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[2] = vo[2] + rcis;\
	vo[3] = vo[3] + rsic;\
	vo[6] = vo[6] + rsic;\
	vo[7] = vo[7] - rcis;\
	vo[10] = vo[10] - rcis;\
	vo[11] = vo[11] - rsic;\
	vo[14] = vo[14] - rsic;\
	vo[15] = vo[15] + rcis;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + ((idx % Ns) << 1)];\
	s1 = factors[(Ns << 3) + ((idx % Ns) << 1) + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[2] = vo[2] + rsic;\
	vo[3] = vo[3] - rcis;\
	vo[4] = vo[4] - rcis;\
	vo[5] = vo[5] - rsic;\
	vo[6] = vo[6] - rsic;\
	vo[7] = vo[7] + rcis;\
	vo[8] = vo[8] + rcis;\
	vo[9] = vo[9] + rsic;\
	vo[10] = vo[10] + rsic;\
	vo[11] = vo[11] - rcis;\
	vo[12] = vo[12] - rcis;\
	vo[13] = vo[13] - rsic;\
	vo[14] = vo[14] - rsic;\
	vo[15] = vo[15] + rcis;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + (idx % Ns) * 3];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 3 + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[12] = vo[12] + rsic;\
	vo[13] = vo[13] - rcis;\
	vo[8] = vo[8] - rcis;\
	vo[9] = vo[9] - rsic;\
	vo[4] = vo[4] - rsic;\
	vo[5] = vo[5] + rcis;\
	c1 = factors[(Ns << 3) + (idx % Ns) * 3 + Ns];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 3 + Ns + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[6] = vo[6] + rcis;\
	vo[7] = vo[7] + rsic;\
	vo[2] = vo[2] + rsic;\
	vo[3] = vo[3] - rcis;\
	vo[14] = vo[14] - rcis;\
	vo[15] = vo[15] - rsic;\
	vo[10] = vo[10] - rsic;\
	vo[11] = vo[11] + rcis;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + ((idx % Ns) << 2)];\
	s1 = factors[(Ns << 3) + ((idx % Ns) << 2) + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[4] = vo[4] + rcis;\
	vo[5] = vo[5] + rsic;\
	vo[8] = vo[8] + rcis;\
	vo[9] = vo[9] + rsic;\
	vo[12] = vo[12] + rcis;\
	vo[13] = vo[13] + rsic;\
	vo[2] = vo[2] - rcis;\
	vo[3] = vo[3] - rsic;\
	vo[6] = vo[6] - rcis;\
	vo[7] = vo[7] - rsic;\
	vo[10] = vo[10] - rcis;\
	vo[11] = vo[11] - rsic;\
	vo[14] = vo[14] - rcis;\
	vo[15] = vo[15] - rsic;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + (idx % Ns) * 5];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 5 + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[4] = vo[4] + rsic;\
	vo[5] = vo[5] - rcis;\
	vo[8] = vo[8] - rcis;\
	vo[9] = vo[9] - rsic;\
	vo[12] = vo[12] - rsic;\
	vo[13] = vo[13] + rcis;\
	c1 = factors[(Ns << 3) + (idx % Ns) * 5 + Ns];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 5 + Ns + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[10] = vo[10] + rcis;\
	vo[11] = vo[11] + rsic;\
	vo[14] = vo[14] + rsic;\
	vo[15] = vo[15] - rcis;\
	vo[2] = vo[2] - rcis;\
	vo[3] = vo[3] - rsic;\
	vo[6] = vo[6] - rsic;\
	vo[7] = vo[7] + rcis;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + (idx % Ns) * 6];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 6 + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[14] = vo[14] + rsic;\
	vo[15] = vo[15] - rcis;\
	vo[12] = vo[12] - rcis;\
	vo[13] = vo[13] - rsic;\
	vo[10] = vo[10] - rsic;\
	vo[11] = vo[11] + rcis;\
	vo[8] = vo[8] + rcis;\
	vo[9] = vo[9] + rsic;\
	vo[6] = vo[6] + rsic;\
	vo[7] = vo[7] - rcis;\
	vo[4] = vo[4] - rcis;\
	vo[5] = vo[5] - rsic;\
	vo[2] = vo[2] - rsic;\
	vo[3] = vo[3] + rcis;\
	\
	idxS += (N>>3);\
	v0 = data0[idxS];\
	v1 = data0[idxS + N];\
	c1 = factors[(Ns << 3) + (idx % Ns) * 7];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 7 + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[0] = vo[0] + rcis;\
	vo[1] = vo[1] + rsic;\
	vo[12] = vo[12] + rsic;\
	vo[13] = vo[13] - rcis;\
	vo[8] = vo[8] - rcis;\
	vo[9] = vo[9] - rsic;\
	vo[4] = vo[4] - rsic;\
	vo[5] = vo[5] + rcis;\
	c1 = factors[(Ns << 3) + (idx % Ns) * 7 + Ns];\
	s1 = factors[(Ns << 3) + (idx % Ns) * 7 + Ns + (NPart << 1)];\
	rcis = v0 * c1 - v1 * s1;\
	rsic = v0 * s1 + v1 * c1;\
	vo[14] = vo[14] + rcis;\
	vo[15] = vo[15] + rsic;\
	vo[10] = vo[10] + rsic;\
	vo[11] = vo[11] - rcis;\
	vo[6] = vo[6] - rcis;\
	vo[7] = vo[7] - rsic;\
	vo[2] = vo[2] - rsic;\
	vo[3] = vo[3] + rcis;\
	\
	idxD = (idx / Ns) * (Ns<<3) + (idx % Ns);\
	data1[idxD] = vo[0];\
	data1[idxD + N] = vo[1];\
	idxD += Ns;\
	data1[idxD] = vo[2];\
	data1[idxD + N] = vo[3];\
	idxD += Ns;\
	data1[idxD] = vo[4];\
	data1[idxD + N] = vo[5];\
	idxD += Ns;\
	data1[idxD] = vo[6];\
	data1[idxD + N] = vo[7];\
	idxD += Ns;\
	data1[idxD] = vo[8];\
	data1[idxD + N] = vo[9];\
	idxD += Ns;\
	data1[idxD] = vo[10];\
	data1[idxD + N] = vo[11];\
	idxD += Ns;\
	data1[idxD] = vo[12];\
	data1[idxD + N] = vo[13];\
	idxD += Ns;\
	data1[idxD] = vo[14];\
	data1[idxD + N] = vo[15];\
}\
";

char* strKernelTwiddle = "\
__kernel void gpuFftTwiddle(__global float* data0, __constant float* factors, int N, int Ns){\
	float c0, s0, c1, s1, c2, s2;\
	float v0, v1;\
	int idx, id_factor0, id_factor1;\
	\
	idx = get_global_id(0);\
	v0 = data0[idx];\
	v1 = data0[idx + N];\
	id_factor0 = idx / Ns;\
	id_factor1 = idx % Ns;\
	id_factor1 = id_factor0 * id_factor1;\
	\
	id_factor0 = id_factor1 / Ns;\
	c0 = factors[Ns + id_factor0];\
	s0 = factors[Ns + (Ns << 1) + id_factor0];\
	id_factor0 = id_factor1 % Ns;\
	c1 = factors[id_factor0];\
	s1 = factors[(Ns << 1) + id_factor0];\
	\
	c2 = c0 * c1 - s0 * s1;\
	s2 = c0 * s1 + c1 * s0;\
	data0[idx] = v0 * c2 - v1 * s2;\
	data0[idx + N] = v0 * s2 + v1 * c2;\
}\
";

char* strKernelPermutation = "\
__kernel void gpuFftPermutation(__global float* data0, __global float* data1, int N, int Ns, int NsConverse){\
	int idx, idin, id_factor0, id_factor1;\
	float v0, v1;\
	\
	idx = get_global_id(0);\
	id_factor0 = idx / Ns;\
	id_factor1 = idx % Ns;\
	idin = id_factor1 * NsConverse + id_factor0;\
	v0 = data0[idin];\
	v1 = data0[idin + N];\
	data1[idx] = v0;\
	data1[idx + N] = v1;\
}\
";

/*
char* strKernelPermutation = "\
__kernel void gpuFftPermutation(__global float* data0, __global float* data1, int N, int Ns, int NsConverse){\
	int idx, id_factor0, id_factor1;\
	float v0, v1;\
	\
	idx = get_global_id(0);\
	v0 = data0[idx];\
	v1 = data0[idx + N];\
	id_factor0 = idx / NsConverse;\
	id_factor1 = idx % NsConverse;\
	idx = id_factor1 * Ns + id_factor0;\
	data1[idx] = v0;\
	data1[idx + N] = v1;\
}\
";*/
